/**
 * Copyright (c) Huawei Technologies Co., Ltd. 2023-2024. All rights reserved.
 *
 * Licensed under the Apache License, Version 2.0 (the "License");
 * you may not use this file except in compliance with the License.
 * You may obtain a copy of the License at
 *
 * http://www.apache.org/licenses/LICENSE-2.0
 *
 * Unless required by applicable law or agreed to in writing, software
 * distributed under the License is distributed on an "AS IS" BASIS,
 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 * See the License for the specific language governing permissions and
 * limitations under the License.
 */

/*!
 * \file flash_attention_score_grad.cpp
 * \brief
 */
#include "kernel_operator.h"

#include "flash_attention_score_grad_empty.h"
#include "flash_attention_score_grad_post.h"
#include "flash_attention_score_grad_us1s2_bbn2gs1s2.h"
#include "flash_attention_score_grad_us1s2_bbn2gs1s2_pre.h"
#include "flash_attention_score_grad_u_s1s2_b_bn2_dropout.h"
#include "flash_attention_score_grad_u_s1s2_b_bn2.h"
#include "flash_attention_score_grad_ngs1s2_bn.h"
#include "flash_attention_score_grad_ngs1s2_bn_dropout.h"
#include "flash_attention_score_grad_bngs1s2_b.h"
#include "flash_attention_score_grad_bngs1s2_b_dropout.h"

constexpr MatmulConfig MM_CFG_EXCEED = GetNormalConfig(true);
constexpr MatmulConfig MM_CFG_NORMAL = GetNormalConfig(false);
constexpr CubeFormat MM_NZ_OUT_FORMAT = CubeFormat::NZ;
constexpr CubeFormat MM_ND_OUT_FORMAT = CubeFormat::ND_ALIGN;
constexpr CubeFormat MM_ND_OUT_NOALIGN = CubeFormat::ND;
constexpr uint64_t INPUT_NONE = 0;
constexpr uint64_t INPUT_EXIST = 1;
constexpr uint32_t INPUT_DISABLE = 0;
constexpr uint32_t INPUT_ENABLE = 1;

constexpr static const uint64_t BSH = 0;
constexpr static const uint64_t SBH = 1;
constexpr static const uint64_t BNSD = 2;
constexpr static const uint64_t BSND = 3;
constexpr static const uint64_t TND = 4;
constexpr static const uint32_t BNGSD = 0;
constexpr static const uint32_t SBNGD = 1;
constexpr static const uint32_t BSNGD = 2;
constexpr static const uint32_t S1S2_TND = 3;

#define INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(INPUT_TYPE, IS_ATTEN_MASK, IS_PSE, IS_DROP, MM_OUT_FORMAT, INPUT_LAYOUT) \
    do {                                                                                                            \
        GET_TILING_DATA_WITH_STRUCT(FlashAttentionScoreGradTilingDataUs1s2Bbn2gs1s2, tiling_data_in, tiling_data);  \
        const FlashAttentionScoreGradTilingDataUs1s2Bbn2gs1s2* __restrict tilingData = &tiling_data_in;             \
        const TCubeTiling* __restrict bmm1tiling = &(tilingData->mm1TilingData);                                    \
        const TCubeTiling* __restrict bmm2tiling = &(tilingData->mm2TilingData);                                    \
        const TCubeTiling* __restrict bmm3tiling = &(tilingData->mm3TilingData);                                    \
        FlashAttentionScoreGradS1S2BNGS1S2Pre<INPUT_TYPE, float, INPUT_LAYOUT> opPre;                                \
        opPre.Init(dq, dk, dv, actual_seq_kvlen, drop_mask, user, tilingData, &pipeIn);                             \
        opPre.Process();                                                                                            \
        opPre.SyncALLCores();                                                                                       \
        pipeIn.Destroy();                                                                                           \
                                                                                                                    \
        TPipe pipeBase;                                                                                             \
        FlashAttentionScoreGradUs1s2Bbn2gs1s2<INPUT_TYPE, float, IS_ATTEN_MASK, IS_PSE, IS_DROP, MM_OUT_FORMAT, INPUT_LAYOUT> op;   \
        REGIST_MATMUL_OBJ(&pipeBase, GetSysWorkSpacePtr(),                                                          \
                        op.mm1, bmm1tiling,                                                                         \
                        op.mm3, bmm2tiling,                                                                         \
                        op.mm4, bmm3tiling);                                                                        \
        op.Init(key, value, dy, query, pse_shift, drop_mask, atten_mask, attention_in, softmax_max, softmax_sum,    \
                actual_seq_qlen, actual_seq_kvlen, dq, dk, dv, dpse, user, tilingData, &pipeBase);                  \
        op.Process();                                                                                               \
        op.SyncALLCores();                                                                                          \
        pipeBase.Destroy();                                                                                         \
        TPipe pipePost;                                                                                             \
        FlashAttentionScoreGradPost<INPUT_TYPE, FlashAttentionScoreGradTilingDataUs1s2Bbn2gs1s2, true> opPost; \
        opPost.Init(dq, dk, dv, user, tilingData, &pipePost);                                                       \
        opPost.Process();                                                                                           \
    } while (0)

#define INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(INPUT_TYPE, MM_CONFIG, CUBE_FORMAT, PSE_CFG, ATTEN_MASK_CFG, DROPOUT_CFG,   \
                                         INPUT_LAYOUT)                                                               \
  do {                                                                                                               \
    GET_TILING_DATA_WITH_STRUCT(FlashAttentionScoreGradTilingDataUs1s2Bbn2, tiling_data_in, tiling_data);            \
    const FlashAttentionScoreGradTilingDataUs1s2Bbn2* __restrict tilingData = &tiling_data_in;                       \
    bool dropBitMode = tilingData->opInfo.S2 % tilingData->splitCoreParams.baseN == 0;                               \
    if (tilingData->opInfo.layout == TND) {                                                                          \
        for (uint32_t i = 0; i < tilingData->opInfo.B; i++) {                                                        \
            uint32_t seq2;                                                                                           \
            if (unlikely(i == 0)) {                                                                                  \
                seq2 = ((__gm__ int64_t *)actual_seq_kvlen)[i];                                                      \
            } else {                                                                                                 \
                seq2 = ((__gm__ int64_t *)actual_seq_kvlen)[i] - ((__gm__ int64_t *)actual_seq_kvlen)[i-1];          \
            }                                                                                                        \
            dropBitMode = (dropBitMode && (seq2 % tilingData->splitCoreParams.baseN == 0));                          \
        }                                                                                                            \
    }                                                                                                                \
    if (g_coreType == AIV && tilingData->opInfo.keepProb < 1 && !dropBitMode) {                                      \
      FAGUs1s2Bbn2Dropout opDropout;                                                                                 \
      opDropout.Init(drop_mask, actual_seq_qlen, actual_seq_kvlen, user, tilingData, &pipeIn);                       \
      opDropout.Process();                                                                                           \
    }                                                                                                                \
    pipeIn.Destroy();                                                                                                \
    TPipe pipeOp;                                                                                                    \
    const TCubeTiling* __restrict bmm1tiling = &(tilingData->mm1TilingData);                                         \
    const TCubeTiling* __restrict bmm31tiling = &(tilingData->mm31TilingData);                                       \
    const TCubeTiling* __restrict bmm4tiling = &(tilingData->mm4TilingData);                                         \
    FlashAttentionScoreGradUs1s2Bbn<INPUT_TYPE, float, MM_CONFIG, CUBE_FORMAT, PSE_CFG, ATTEN_MASK_CFG, DROPOUT_CFG, \
        INPUT_LAYOUT> op;                                                                                            \
    REGIST_MATMUL_OBJ(&pipeOp, GetSysWorkSpacePtr(), op.mm1, bmm1tiling, op.mm4, bmm4tiling,                         \
                      op.mm3_1, bmm31tiling);                                                                        \
    op.Init(query, key, value, dy, pse_shift, drop_mask, padding_mask, atten_mask, softmax_max, softmax_sum,         \
            softmax_in, actual_seq_qlen, actual_seq_kvlen, attention_in, dq, dk, dv, dpse, user, tilingData,         \
            &pipeOp);                                                                                                \
    op.Process();                                                                                                    \
    pipeOp.Destroy();                                                                                                \
    TPipe pipeCast;                                                                                                  \
    FlashAttentionScoreGradPost<INPUT_TYPE, FlashAttentionScoreGradTilingDataUs1s2Bbn2, true> opCast;      \
    opCast.Init(dq, dk, dv, user, tilingData, &pipeCast);                                                            \
    opCast.Process();                                                                                                \
  } while (0)

#define INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(INPUT_TYPE, MM_CONFIG, CUBE_FORMAT, PSE_CFG, ATTEN_MASK_CFG,     \
                                                    DROPOUT_CFG, INPUT_LAYOUT)                                       \
  do {                                                                                                               \
    GET_TILING_DATA_WITH_STRUCT(FlashAttentionScoreGradTilingDataUs1s2Bbn2, tiling_data_in, tiling_data);            \
    const FlashAttentionScoreGradTilingDataUs1s2Bbn2* __restrict tilingData = &tiling_data_in;                       \
    const TCubeTiling* __restrict bmm1tiling = &(tilingData->mm1TilingData);                                         \
    const TCubeTiling* __restrict bmm31tiling = &(tilingData->mm31TilingData);                                       \
    const TCubeTiling* __restrict bmm4tiling = &(tilingData->mm4TilingData);                                         \
    FlashAttentionScoreGradUs1s2Bbn<INPUT_TYPE, float, MM_CONFIG, CUBE_FORMAT, PSE_CFG, ATTEN_MASK_CFG, DROPOUT_CFG, \
        INPUT_LAYOUT> op;                                                                                            \
    REGIST_MATMUL_OBJ(&pipeIn, GetSysWorkSpacePtr(), op.mm1, bmm1tiling, op.mm4, bmm4tiling,                         \
                      op.mm3_1, bmm31tiling);                                                                        \
    op.Init(query, key, value, dy, pse_shift, drop_mask, padding_mask, atten_mask, softmax_max, softmax_sum,         \
            softmax_in, actual_seq_qlen, actual_seq_kvlen, attention_in, dq, dk, dv, dpse, user, tilingData,         \
            &pipeIn);                                                                                                \
    op.Process();                                                                                                    \
    pipeIn.Destroy();                                                                                                \
    TPipe pipeCast;                                                                                                  \
    FlashAttentionScoreGradPost<INPUT_TYPE, FlashAttentionScoreGradTilingDataUs1s2Bbn2, true> opCast;     \
    opCast.Init(dq, dk, dv, user, tilingData, &pipeCast);                                                            \
    opCast.Process();                                                                                                \
  } while (0)

#define INVOKE_FAG_GENERAL_S1S2_BNGS1S2_B_IMPL(INPUT_TYPE, layout, MM_CONFIG)                                        \
  do {                                                                                                               \
    GET_TILING_DATA_WITH_STRUCT(FlashAttentionScoreGradUbngs1s2BbTilingData, tiling_data_in, tiling_data);           \
    const FlashAttentionScoreGradUbngs1s2BbTilingData* __restrict tilingData = &tiling_data_in;                      \
    if (g_coreType == AIV && tilingData->opInfo.keepProb < 1 && (tilingData->opInfo.sKV % DROPOUT4BIT_LEN != 0)) {   \
        TPipe pipeInDrop;                                                                                            \
        FAGNgs1s2B2Dropout opDropout;                                                                                \
        opDropout.Init(drop_mask, user, &tiling_data_in, &pipeInDrop);                                               \
        opDropout.Process();                                                                                         \
        pipeInDrop.Destroy();                                                                                        \
    }                                                                                                                \
    pipeIn.Destroy();                                                                                                \
    TPipe pipeOp;                                                                                                    \
    const TCubeTiling* __restrict bmm1tiling = &(tilingData->mm1AndMm2TilingData);                                   \
    const TCubeTiling* __restrict bmm2tiling = &(tilingData->mm1AndMm2TilingData);                                   \
    const TCubeTiling* __restrict bmm3tiling = &(tilingData->mm31TilingData);                                        \
    const TCubeTiling* __restrict bmm4tiling = &(tilingData->mm32AndMm4TilingData);                                  \
    FlashAttentionScoreGradUngs1s2Bb<INPUT_TYPE, float, MM_CONFIG, layout> op;                                       \
    REGIST_MATMUL_OBJ(&pipeOp, GetSysWorkSpacePtr(), op.mm1, bmm1tiling, op.mm2, bmm2tiling, op.mm31, bmm3tiling,    \
                      op.mm32, bmm4tiling, op.mm4, bmm4tiling);                                                      \
    op.Init(key, value, dy, query, pse_shift, drop_mask, atten_mask, attention_in, softmax_max, softmax_sum, dq, dk, \
            dv, user, tilingData, &pipeOp);                                                                          \
    op.Process();                                                                                                    \
    pipeOp.Destroy();                                                                                                \
    TPipe pipeMuls;                                                                                                  \
    FlashAttentionScoreGradPost<INPUT_TYPE, FlashAttentionScoreGradUbngs1s2BbTilingData, false> opMuls;    \
    opMuls.Init(dq, dk, dv, user, tilingData, &pipeMuls);                                                                \
    opMuls.Process();                                                                                                \
    pipeMuls.Destroy();                                                                                              \
  } while (0)

#define INVOKE_FAG_GENERAL_S1S2_BNGS1S2_IMPL(INPUT_TYPE, layout, MM_CONFIG)                                          \
  do {                                                                                                               \
    GET_TILING_DATA_WITH_STRUCT(FlashAttentionScoreGradTilingDataUngs1s2Bbn, tiling_data_in, tiling_data);           \
    const FlashAttentionScoreGradTilingDataUngs1s2Bbn* __restrict tilingData = &tiling_data_in;                      \
    if (g_coreType == AIV && tilingData->opInfo.keepProb < 1 &&                                                      \
        (tilingData->opInfo.sKV % BN_DROPOUT4BIT_LEN != 0)) {                                                        \
        TPipe pipeInDp;                                                                                              \
        FAGNgs1s2Bn2Dropout opDropout;                                                                               \
        opDropout.Init(drop_mask, user, &tiling_data_in, &pipeInDp);                                                 \
        opDropout.Process();                                                                                         \
        pipeInDp.Destroy();                                                                                          \
    }                                                                                                                \
    pipeIn.Destroy();                                                                                                \
    TPipe pipeOp;                                                                                                    \
    const TCubeTiling* __restrict bmm1tiling = &(tilingData->mm1AndMm2TilingData);                                   \
    const TCubeTiling* __restrict bmm2tiling = &(tilingData->mm1AndMm2TilingData);                                   \
    const TCubeTiling* __restrict bmm3tiling = &(tilingData->mm31TilingData);                                        \
    const TCubeTiling* __restrict bmm4tiling = &(tilingData->mm32AndMm4TilingData);                                  \
    FlashAttentionScoreGradUngs1s2Bbn<INPUT_TYPE, float, MM_CONFIG, true, layout> op;                                \
    REGIST_MATMUL_OBJ(&pipeOp, GetSysWorkSpacePtr(), op.mm1, bmm1tiling, op.mm2, bmm2tiling, op.mm31,                \
                      bmm3tiling, op.mm32, bmm4tiling, op.mm4, bmm4tiling);                                          \
    op.Init(key, value, dy, query, pse_shift, drop_mask, atten_mask, attention_in, softmax_max, softmax_sum,         \
            dq, dk, dv, user, tilingData, &pipeOp);                                                                  \
    op.Process();                                                                                                    \
    pipeOp.Destroy();                                                                                                \
    TPipe pipeMuls;                                                                                                  \
    FlashAttentionScoreGradPost<INPUT_TYPE, FlashAttentionScoreGradTilingDataUngs1s2Bbn, false> opMuls;    \
    opMuls.Init(dq, dk, dv, user, tilingData, &pipeMuls);                                                                \
    opMuls.Process();                                                                                                \
    pipeMuls.Destroy();                                                                                              \
  } while (0)

// implementation of kernel function
extern "C" __global__ __aicore__ void flash_attention_score_grad(
    __gm__ uint8_t* query, __gm__ uint8_t* key, __gm__ uint8_t* value, __gm__ uint8_t* dy, __gm__ uint8_t* pse_shift,
    __gm__ uint8_t* drop_mask, __gm__ uint8_t* padding_mask, __gm__ uint8_t* atten_mask, __gm__ uint8_t* softmax_max,
    __gm__ uint8_t* softmax_sum, __gm__ uint8_t* softmax_in, __gm__ uint8_t* attention_in, __gm__ uint8_t* prefix,
    __gm__ uint8_t* actual_seq_qlen, __gm__ uint8_t* actual_seq_kvlen, __gm__ uint8_t* dq, __gm__ uint8_t* dk,
    __gm__ uint8_t* dv, __gm__ uint8_t* dpse, __gm__ uint8_t* workspace, __gm__ uint8_t* tiling_data) {
  TPipe pipeIn;
  set_mask_norm();
  __gm__ uint8_t* user = GetUserWorkspace(workspace);

// --------------------------------------------float16 tilingkey------------------------------------------------------
#if (ORIG_DTYPE_QUERY == DT_FLOAT16)
  // -----------------------1.1 start---------------------------------
  if (TILING_KEY_IS(10000000000111003434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000000011003434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000000101003434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000000001003434UL)) {
    // attention_mask:0, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000000110003434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000000010003434UL)) {
    // attention_mask:0, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000000100003434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000000000003434UL)) {
    // attention_mask:0, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001111003434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001011003434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001101003434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001001003434UL)) {
    // attention_mask:0, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001110003434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001010003434UL)) {
    // attention_mask:0, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001100003434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001000003434UL)) {
    // attention_mask:0, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  // 1
  } else if (TILING_KEY_IS(10000000000111013434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000000011013434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000000101013434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000000001013434UL)) {
    // attention_mask:0, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000000110013434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000000010013434UL)) {
    // attention_mask:0, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000000100013434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000000000013434UL)) {
    // attention_mask:0, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001111013434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001011013434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001101013434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001001013434UL)) {
    // attention_mask:0, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001110013434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001010013434UL)) {
    // attention_mask:0, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001100013434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001000013434UL)) {
    // attention_mask:0, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  // 2
  } else if (TILING_KEY_IS(10000000000111023434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000000011023434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000000101023434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000000001023434UL)) {
    // attention_mask:0, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000000110023434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000000010023434UL)) {
    // attention_mask:0, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000000100023434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000000000023434UL)) {
    // attention_mask:0, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001111023434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001011023434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001101023434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001001023434UL)) {
    // attention_mask:0, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001110023434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001010023434UL)) {
    // attention_mask:0, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001100023434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001000023434UL)) {
    // attention_mask:0, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  // 3
  } else if (TILING_KEY_IS(10000000000111033434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000000011033434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000000101033434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000000001033434UL)) {
    // attention_mask:0, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000000110033434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000000010033434UL)) {
    // attention_mask:0, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000000100033434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000000000033434UL)) {
    // attention_mask:0, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001111033434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001011033434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001101033434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001001033434UL)) {
    // attention_mask:0, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001110033434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001010033434UL)) {
    // attention_mask:0, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001100033434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001000033434UL)) {
    // attention_mask:0, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(half, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  // -----------------------1.1 end---------------------------------

  // -----------------------1.2 start---------------------------------
  // For BSND
    // pse atten_mask dropout 均不存在
  } else if (TILING_KEY_IS(10000000000000000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000000010000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000000001000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000000011000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
    // pse单独存在
  } else if (TILING_KEY_IS(10000000000100000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000000110000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000000101000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000000111000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
    // atten_mask单独存在
  } else if (TILING_KEY_IS(10000000001000000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000001010000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000001001000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000001011000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
    // dropout单独存在
  } else if (TILING_KEY_IS(10000000010000000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000010010000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000010001000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000010011000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BSND);
    return;
    // pse atten_mask存在
  } else if (TILING_KEY_IS(10000000001100000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000001110000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000001101000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000001111000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
    // pse dropout存在
  } else if (TILING_KEY_IS(10000000010100000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000010110000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000010101000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000010111000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BSND);
    return;
    // atten_mask dropout存在
  } else if (TILING_KEY_IS(10000000011000000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000011010000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000011001000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000011011000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BSND);
    return;
    // 均存在
  } else if (TILING_KEY_IS(10000000011100000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000011110000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000011101000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000011111000134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BSND);
    return;
  // For SBH
    // pse atten_mask dropout 均不存在
  } else if (TILING_KEY_IS(10000000000000010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000000010010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000000001010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000000011010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
    // pse单独存在
  } else if (TILING_KEY_IS(10000000000100010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000000110010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000000101010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000000111010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
    // atten_mask单独存在
  } else if (TILING_KEY_IS(10000000001000010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000001010010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000001001010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000001011010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
    // dropout单独存在
  } else if (TILING_KEY_IS(10000000010000010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000010010010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000010001010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000010011010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, SBH);
    return;
    // pse atten_mask存在
  } else if (TILING_KEY_IS(10000000001100010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000001110010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000001101010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000001111010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
    // pse dropout存在
  } else if (TILING_KEY_IS(10000000010100010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000010110010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000010101010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000010111010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, SBH);
    return;
    // atten_mask dropout存在
  } else if (TILING_KEY_IS(10000000011000010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000011010010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000011001010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000011011010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, SBH);
    return;
    // 均存在
  } else if (TILING_KEY_IS(10000000011100010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000011110010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000011101010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000011111010134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, SBH);
    return;
  // For BNSD
  // pse atten_mask dropout 均不存在
  } else if (TILING_KEY_IS(10000000000000020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000000010020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000000001020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000000011020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
    // pse单独存在
  } else if (TILING_KEY_IS(10000000000100020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000000110020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000000101020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000000111020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
    // atten_mask单独存在
  } else if (TILING_KEY_IS(10000000001000020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000001010020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000001001020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000001011020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
    // dropout单独存在
  } else if (TILING_KEY_IS(10000000010000020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000010010020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000010001020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000010011020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
    // pse atten_mask存在
  } else if (TILING_KEY_IS(10000000001100020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000001110020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000001101020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000001111020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
    // pse dropout存在
  } else if (TILING_KEY_IS(10000000010100020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000010110020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000010101020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000010111020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
    // atten_mask dropout存在
  } else if (TILING_KEY_IS(10000000011000020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000011010020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000011001020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000011011020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BNSD);
    return;
    // 均存在
  } else if (TILING_KEY_IS(10000000011100020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000011110020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000011101020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000011111020134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BNSD);
    return;
    // for TND
  } else if (TILING_KEY_IS(10000000000000030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000000010030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000000001030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000000011030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
    // pse单独存在
  } else if (TILING_KEY_IS(10000000000100030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000000110030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000000101030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000000111030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
    // atten_mask单独存在
  } else if (TILING_KEY_IS(10000000001000030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000001010030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000001001030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000001011030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
    // dropout单独存在
  } else if (TILING_KEY_IS(10000000010000030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000010010030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000010001030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000010011030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, TND);
    return;
    // pse atten_mask存在
  } else if (TILING_KEY_IS(10000000001100030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000001110030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000001101030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000001111030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
    // pse dropout存在
  } else if (TILING_KEY_IS(10000000010100030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000010110030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000010101030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000010111030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, TND);
    return;
    // atten_mask dropout存在
  } else if (TILING_KEY_IS(10000000011000030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000011010030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000011001030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000011011030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, TND);
    return;
    // 均存在
  } else if (TILING_KEY_IS(10000000011100030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000011110030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000011101030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000011111030134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(half, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, TND);
    return;
  // -----------------------1.2 end---------------------------------

  } else if (TILING_KEY_IS(10000000000000003199UL)) { // BSH BSND & FLOAT16_PRECISION
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_IMPL(half, LayoutMode::BSNGD, MM_CFG_NORMAL);
    return;
  } else if (TILING_KEY_IS(10000000000000013199UL)) { // SBH & FLOAT16_PRECISION
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_IMPL(half, LayoutMode::SBNGD, MM_CFG_NORMAL);
    return;
  } else if (TILING_KEY_IS(10000000000000023199UL)) { // BNSD & FLOAT16_PRECISION
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_IMPL(half, LayoutMode::BNGS1S2, MM_CFG_NORMAL);
    return;
  } else if (TILING_KEY_IS(10000000000001013199)) { // SBH & FLOAT16_PRECISION & speical MM tilingkey
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_IMPL(half, LayoutMode::SBNGD, MM_CFG_EXCEED);
    return;
  } else if (TILING_KEY_IS(10000000000000103099UL)) {
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_B_IMPL(half, LayoutMode::BSNGD, MM_CFG_NORMAL);
    return;
  } else if (TILING_KEY_IS(10000000000000113099UL)) {
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_B_IMPL(half, LayoutMode::SBNGD, MM_CFG_NORMAL);
    return;
  } else if (TILING_KEY_IS(10000000000001113099UL)) {
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_B_IMPL(half, LayoutMode::SBNGD, MM_CFG_EXCEED);
    return;
  } else if (TILING_KEY_IS(10000000000000123099UL)) {
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_B_IMPL(half, LayoutMode::BNGS1S2, MM_CFG_NORMAL);
    return;
  } else if (TILING_KEY_IS(10000000000000133099UL)) {
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_B_IMPL(half, LayoutMode::BSNGD, MM_CFG_NORMAL);
    return;
  }
#endif
// --------------------------------------------------------------------------------------------------------------------

// --------------------------------------------bfloat16 tilingkey------------------------------------------------------
#if (ORIG_DTYPE_QUERY == DT_BF16)
  // -----------------------1.1 start---------------------------------
  if (TILING_KEY_IS(10000000000111002434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000000011002434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000000101002434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000000001002434UL)) {
    // // attention_mask:0, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000000110002434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000000010002434UL)) {
    //  attention_mask:0, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000000100002434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000000000002434UL)) {
    //  attention_mask:0, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001111002434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001011002434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001101002434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001001002434UL)) {
    // // attention_mask:0, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001110002434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001010002434UL)) {
    //  attention_mask:0, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001100002434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  } else if (TILING_KEY_IS(10000000001000002434UL)) {
    //  attention_mask:0, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BSNGD);
    return;
  // 1
  } else if (TILING_KEY_IS(10000000000111012434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000000011012434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000000101012434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000000001012434UL)) {
    // // attention_mask:0, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000000110012434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000000010012434UL)) {
    //  attention_mask:0, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000000100012434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000000000012434UL)) {
    //  attention_mask:0, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001111012434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001011012434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001101012434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001001012434UL)) {
    // // attention_mask:0, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001110012434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001010012434UL)) {
    //  attention_mask:0, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001100012434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  } else if (TILING_KEY_IS(10000000001000012434UL)) {
    //  attention_mask:0, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, SBNGD);
    return;
  // 2
  } else if (TILING_KEY_IS(10000000000111022434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000000011022434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000000101022434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000000001022434UL)) {
    // // attention_mask:0, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000000110022434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000000010022434UL)) {
    //  attention_mask:0, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000000100022434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000000000022434UL)) {
    //  attention_mask:0, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001111022434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001011022434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001101022434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001001022434UL)) {
    // // attention_mask:0, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001110022434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001010022434UL)) {
    //  attention_mask:0, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001100022434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  } else if (TILING_KEY_IS(10000000001000022434UL)) {
    //  attention_mask:0, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, BNGSD);
    return;
  // 3
  } else if (TILING_KEY_IS(10000000000111032434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000000011032434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000000101032434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000000001032434UL)) {
    // // attention_mask:0, pse:0, drop:1, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000000110032434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000000010032434UL)) {
    //  attention_mask:0, pse:1, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000000100032434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000000000032434UL)) {
    //  attention_mask:0, pse:0, drop:0, mm_out:nd
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_ND_OUT_NOALIGN, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001111032434UL)) {
    // attention_mask:1, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001011032434UL)) {
    // attention_mask:0, pse:1, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001101032434UL)) {
    // attention_mask:1, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001001032434UL)) {
    // // attention_mask:0, pse:0, drop:1, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_ENABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001110032434UL)) {
    // attention_mask:1, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001010032434UL)) {
    //  attention_mask:0, pse:1, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_ENABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001100032434UL)) {
    // attention_mask:1, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_ENABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  } else if (TILING_KEY_IS(10000000001000032434UL)) {
    //  attention_mask:0, pse:0, drop:0, mm_out:nz
    INVOKE_FAG_GENERAL_S1S2_BN2GS1S2_IMPL(bfloat16_t, INPUT_DISABLE, INPUT_DISABLE, INPUT_DISABLE, MM_NZ_OUT_FORMAT, S1S2_TND);
    return;
  // -----------------------1.1 end---------------------------------

  // -----------------------1.2 start-------------------------------
  // For BSND
  } else if (TILING_KEY_IS(10000000000000002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000000010002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000000001002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000000011002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
    // pse单独存在
  } else if (TILING_KEY_IS(10000000000100002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000000110002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000000101002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000000111002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BSND);
    return;
    // atten_mask单独存在
  } else if (TILING_KEY_IS(10000000001000002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000001010002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000001001002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000001011002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
    // dropout单独存在
  } else if (TILING_KEY_IS(10000000010000002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000010010002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000010001002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000010011002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BSND);
    return;
    // pse atten_mask存在
  } else if (TILING_KEY_IS(10000000001100002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000001110002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000001101002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
  } else if (TILING_KEY_IS(10000000001111002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BSND);
    return;
    // pse dropout存在
  } else if (TILING_KEY_IS(10000000010100002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000010110002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000010101002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000010111002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BSND);
    return;
    // atten_mask dropout存在
  } else if (TILING_KEY_IS(10000000011000002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000011010002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000011001002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000011011002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BSND);
    return;
    // 均存在
  } else if (TILING_KEY_IS(10000000011100002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000011110002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000011101002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BSND);
    return;
  } else if (TILING_KEY_IS(10000000011111002134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BSND);
    return;
  // For SBH
  } else if (TILING_KEY_IS(10000000000000012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000000010012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000000001012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000000011012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
    // pse单独存在
  } else if (TILING_KEY_IS(10000000000100012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000000110012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000000101012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000000111012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, SBH);
    return;
    // atten_mask单独存在
  } else if (TILING_KEY_IS(10000000001000012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000001010012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000001001012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000001011012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
    // dropout单独存在
  } else if (TILING_KEY_IS(10000000010000012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000010010012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000010001012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000010011012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, SBH);
    return;
    // pse atten_mask存在
  } else if (TILING_KEY_IS(10000000001100012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000001110012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000001101012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
  } else if (TILING_KEY_IS(10000000001111012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, SBH);
    return;
    // pse dropout存在
  } else if (TILING_KEY_IS(10000000010100012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000010110012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000010101012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000010111012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, SBH);
    return;
    // atten_mask dropout存在
  } else if (TILING_KEY_IS(10000000011000012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000011010012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000011001012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000011011012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, SBH);
    return;
    // 均存在
  } else if (TILING_KEY_IS(10000000011100012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000011110012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000011101012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, SBH);
    return;
  } else if (TILING_KEY_IS(10000000011111012134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, SBH);
    return;
  // For BNSD
  // pse atten_mask dropout 均不存在
  } else if (TILING_KEY_IS(10000000000000022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000000010022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000000001022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000000011022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
    // pse单独存在
  } else if (TILING_KEY_IS(10000000000100022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000000110022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000000101022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000000111022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, BNSD);
    return;
    // atten_mask单独存在
  } else if (TILING_KEY_IS(10000000001000022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000001010022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000001001022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000001011022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
    // dropout单独存在
  } else if (TILING_KEY_IS(10000000010000022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000010010022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000010001022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000010011022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
    // pse atten_mask存在
  } else if (TILING_KEY_IS(10000000001100022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000001110022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000001101022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000001111022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, BNSD);
    return;
    // pse dropout存在
  } else if (TILING_KEY_IS(10000000010100022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000010110022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000010101022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000010111022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, BNSD);
    return;
    // atten_mask dropout存在
  } else if (TILING_KEY_IS(10000000011000022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000011010022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000011001022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000011011022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, BNSD);
    return;
    // 均存在
  } else if (TILING_KEY_IS(10000000011100022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000011110022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000011101022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BNSD);
    return;
  } else if (TILING_KEY_IS(10000000011111022134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, BNSD);
    return;
    // for TND
  } else if (TILING_KEY_IS(10000000000000032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000000010032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000000001032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000000011032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
    // pse单独存在
  } else if (TILING_KEY_IS(10000000000100032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000000110032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000000101032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000000111032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE,
                                                INPUT_NONE, TND);
    return;
    // atten_mask单独存在
  } else if (TILING_KEY_IS(10000000001000032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000001010032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000001001032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000001011032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
    // dropout单独存在
  } else if (TILING_KEY_IS(10000000010000032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000010010032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000010001032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000010011032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_NONE, INPUT_EXIST, TND);
    return;
    // pse atten_mask存在
  } else if (TILING_KEY_IS(10000000001100032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000001110032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000001101032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
  } else if (TILING_KEY_IS(10000000001111032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_NO_DROPOUT_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                                INPUT_NONE, TND);
    return;
    // pse dropout存在
  } else if (TILING_KEY_IS(10000000010100032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000010110032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000010101032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000010111032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_NONE, INPUT_EXIST, TND);
    return;
    // atten_mask dropout存在
  } else if (TILING_KEY_IS(10000000011000032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000011010032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000011001032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000011011032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_NONE, INPUT_EXIST, INPUT_EXIST, TND);
    return;
    // 均存在
  } else if (TILING_KEY_IS(10000000011100032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000011110032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_NORMAL, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000011101032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_ND_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, TND);
    return;
  } else if (TILING_KEY_IS(10000000011111032134UL)) {
    INVOKE_FAG_GENERAL_S1S2_BN2_IMPL(bfloat16_t, MM_CFG_EXCEED, MM_NZ_OUT_FORMAT, INPUT_EXIST, INPUT_EXIST,
                                     INPUT_EXIST, TND);
    return;
  // -----------------------1.2 end---------------------------------

  } else if (TILING_KEY_IS(10000000000000002199UL)) { // BSH BSND & BFLOAT16
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_IMPL(bfloat16_t, LayoutMode::BSNGD, MM_CFG_NORMAL);
    return;
  } else if (TILING_KEY_IS(10000000000000012199UL)) { // SBH & BFLOAT16
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_IMPL(bfloat16_t, LayoutMode::SBNGD, MM_CFG_NORMAL);
    return;
  } else if (TILING_KEY_IS(10000000000000022199UL)) { // BNSD & BFLOAT16
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_IMPL(bfloat16_t, LayoutMode::BNGS1S2, MM_CFG_NORMAL);
    return;
  } else if (TILING_KEY_IS(10000000000001012199UL)) { // SBH & BFLOAT16  & speical MM tilingkey
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_IMPL(bfloat16_t, LayoutMode::SBNGD, MM_CFG_EXCEED);
    return;
  } else if (TILING_KEY_IS(10000000000000102099UL)) {
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_B_IMPL(bfloat16_t, LayoutMode::BSNGD, MM_CFG_NORMAL);
    return;
  } else if (TILING_KEY_IS(10000000000000112099UL)) {
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_B_IMPL(bfloat16_t, LayoutMode::SBNGD, MM_CFG_NORMAL);
    return;
  } else if (TILING_KEY_IS(10000000000001112099UL)) {
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_B_IMPL(bfloat16_t, LayoutMode::SBNGD, MM_CFG_EXCEED);
    return;
  } else if (TILING_KEY_IS(10000000000000122099UL)) {
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_B_IMPL(bfloat16_t, LayoutMode::BNGS1S2, MM_CFG_NORMAL);
    return;
  } else if (TILING_KEY_IS(10000000000000132099UL)) {
    INVOKE_FAG_GENERAL_S1S2_BNGS1S2_B_IMPL(bfloat16_t, LayoutMode::BSNGD, MM_CFG_NORMAL);
    return;
  }
#endif
// --------------------------------------------------------------------------------------------------------------------

  GET_TILING_DATA_WITH_STRUCT(FlashAttentionScoreGradTilingData, tiling_data_in, tiling_data);
  const FlashAttentionScoreGradTilingData* __restrict tilingData = &tiling_data_in;

  if (TILING_KEY_IS(90)) {
    FlashAttentionScoreGradEmpty<DTYPE_DQ> op;
    op.Init(dq, dk, dv, dpse, tilingData);
    op.Process();
  }
}